Merging NCCL 2.21.5-1 #1271
Closed
Merging NCCL 2.21.5-1 #1271
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Add support for IB SHARP 1PPN operation with user buffers. Improve support for MNNVL, add NVLS support and multi-clique support. * Detect the NVLS clique through NVML * Exchange XML between peers in the same NVLS clique and fuse XMLs before creating the topology graph. * Rework bootstrap allgather algorithms to allow for large allgather operations intra-node (XML exchange). Net/IB: add support for dynamic GID detection. * Automatically select RoCEv2/IPv4 interface by default. Allow to select IPv6 or even the network/mask. Reduce NVLS memory usage. * Add stepSize as property of a connection to allow for different sizes on different peers; set it to 128K for NVLink SHARP. Improve tuner loading * Look for more paths, be more consistent with the network device plugin. * Also search for tuner support inside the net plugin. Improve tuner API * Add context to support multi-device per process. Add magic number around comm object to detect comm corruption. * Add some basic check around communicators so that we can report a problem when a communicator gets corrupted or a wrong comm pointer is passed to NCCL. Fix net/IB error path. Github PR ROCm#1164 Fix collnet rail mapping with split comm. Fix packet reordering issue causing bootstrap mismatch * Use a different tag in ncclTransportP2pSetup for the connectInfo exchange and the following barrier. Fix hang when crossNic is inconsistent between ranks. Fix minCompCap/maxCompCap computation. Github issue ROCm#1184
corey-derochie-amd
force-pushed
the
feature/nccl-2.21-sync
branch
2 times, most recently
from
ec54e8f to
059f28c
Compare
Comment on lines
-98
to
+97
|
|
||
| #define NCCL_TOPO_MAX_LINKS 128 |
There was a problem hiding this comment.
RCCL expanded this from 32 to 64 for CPX mode, but NCCL expands all the way to 128, so we've taken the larger value.
Comment on lines
+202
to
+205
| int64_t netId; | ||
| if (connIndex == NCCL_CONN_IDX_P2P_NET) NCCLCHECK(ncclTopoGetIntraNetDev(comm->topo, myInfo->rank, graph, channelId, 1, &netId, &req.netDev)); | ||
| if (req.netDev < 0) NCCLCHECK(ncclTopoGetNetDev(comm, myInfo->rank, graph, channelId, peerInfo->rank, &netId, &req.netDev, &proxyRank)); | ||
| NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netId, 1, &req.useGdr)); |
There was a problem hiding this comment.
A few places use ncclTopoGetIntraNetDev in RCCL to circumvent ncclTopoGetNetDev, so with the netId changes, this logic had to be updated as well.
| @@ -511,16 +538,16 @@ ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* s | |||
| } | |||
| for (int s=0; s<xmlCpu->nSubs; s++) { | |||
| struct ncclXmlNode* node = xmlCpu->subs[s]; | |||
| if (strcmp(node->name, "pci") == 0) NCCLCHECK(ncclTopoAddPci(node, system, cpu)); | |||
| if (strcmp(node->name, "pci") == 0) NCCLCHECK(ncclTopoAddPci(node, system, cpu, systemId)); | |||
There was a problem hiding this comment.
@gilbertlee-amd This looks like a schema change to the XML. Do our models need to be redone? Does the Topology Visualizer need to be updated?
Comment on lines
+2007
to
+2011
| INFO(NCCL_INIT,"ncclCommSplit comm %p rank %d nranks %d cudaDev %d busId %lx parent %p color %d key %d commId 0x%llx - Init START", | ||
| comm, comm->rank, comm->nRanks, comm->cudaDev, comm->busId, job->parent, job->color, job->key, (unsigned long long)hashUniqueId(job->commId)); | ||
| } else { | ||
| INFO(NCCL_INIT,"ncclCommInitRank comm %p rank %d nranks %d cudaDev %d busId %lx commId 0x%llx - Init START", | ||
| comm, comm->rank, comm->nRanks, comm->cudaDev, comm->busId, (unsigned long long)hashUniqueId(job->commId)); |
There was a problem hiding this comment.
NCCL uses nvlsDev instead of cudaDev, but RCCL contains this switch so these NCCL changes were modified to match. This pattern appears elsewhere.
corey-derochie-amd
force-pushed
the
feature/nccl-2.21-sync
branch
from
e68adad to
0e657bc
Compare
corey-derochie-amd
requested review from
gilbertlee-amd,
akolliasAMD,
edgargabriel,
PedramAlizadeh,
nusislam,
nileshnegi,
KawtharShafie,
AtlantaPepsi,
mberenjk and
haripriya-amd
as code owners
Comment on lines
+289
to
+290
| if (xml->maxIndex == xml->maxNodes) { | ||
| WARN("Error : too many XML nodes (max %d)", xml->maxNodes); |
There was a problem hiding this comment.
Since MAX_NODES is no longer set in code, it must be added to the XML for CPX mode.
Comment on lines
+202
to
+203
| #define NCCL_TOPO_XML_MAX_NODES 256 | ||
| #define NCCL_GRAPH_XML_MAX_NODES 4096 |
There was a problem hiding this comment.
@nusislam How do we resolve these against the CPX changes?
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Details
Do not mention proprietary info or link to internal work items in this PR.
Work item: "Internal", or link to GitHub issue (if applicable).
What were the changes?
Merged NCCL v2.21.5-1 into RCCL develop.
Why were the changes made?
Must remain up-to-date with latest NCCL versions.
How was the outcome achieved?
Performed a git-merge, then resolved merge conflicts, then resolved compiler errors, then resolved logical conflicts due to functional changes in NCCL.
Additional Documentation:
NCCL Documentation:
Approval Checklist
Do not approve until these items are satisfied.